
#include "cudaBindings.h"

__global__ void cudaKernel(Sphere *spheres, float *depths, size_t count)
{
	int n = blockIdx.x * blockDim.x + threadIdx.x;
	if (n < count)
	{
		float res = 0.0f;
		for (size_t i = 0; i < count; i++)
			if (i != n)
			{
				res += Sphere::findIntersectionDepth(spheres[i], spheres[n]);
			}
		depths[n] = res / 2.0f;
		__syncthreads();
	}
}

void cudaIntersectionFinder::processKernel(Sphere *spheres, float *depths, size_t count)
{
	//Copying spheres and depths to GPU memory.
	Sphere *spheres_on_gpu;
	float *depths_on_gpu;
	SAFE_CALL(cudaMalloc((void **)&spheres_on_gpu, sizeof(Sphere) * count));
	SAFE_CALL(cudaMemcpy(spheres_on_gpu, spheres, sizeof(Sphere) * count, cudaMemcpyHostToDevice));
	SAFE_CALL(cudaMalloc((void **)&depths_on_gpu, sizeof(float) * count));
	SAFE_CALL(cudaMemcpy(depths_on_gpu, depths, sizeof(float) * count, cudaMemcpyHostToDevice));
	//Launching CUDA-kernel.
	dim3 threads(256);
	dim3 grid(count / 256 + 1);
	cudaKernel<<<threads, grid>>>(spheres_on_gpu, depths_on_gpu, count);
	CHECK_ERROR("CUDA kernel failed");
	//Copying result to host memory and freeing GPU memory.
	SAFE_CALL(cudaMemcpy(depths, depths_on_gpu, sizeof(float) * count, cudaMemcpyDeviceToHost));
	SAFE_CALL(cudaFree(spheres_on_gpu));
	SAFE_CALL(cudaFree(depths_on_gpu));
}
